#include<stdio.h>
#include<cuda.h>
#include<cuda_runtime.h>

const int block_size_x = 16;
const int block_size_y = 16;

const int height = 5;
const int width = 5;

int div_round_up(int ptr, int b) {
	return ((ptr % b) != 0) ? (ptr / b + 1) : (ptr / b);
}

__global__ void kernel_malloc_2d(float *dev_p, size_t pitch, int width, int height)
{
	int tid_x = blockIdx.x * blockDim.x + threadIdx.x;
	int tid_y = blockIdx.y * blockDim.y + threadIdx.y;

	if ((tid_x < width) && (tid_y < height))
	{
		float *row_a = (float *)((char*)dev_p + tid_y * pitch);
		row_a[tid_x] = row_a[tid_x] + 100.0f;
	}
}

int main()
{
	float h_vec[height][width];
	float *dev_p;
	size_t pitch;

	for (int i = 0; i < height; i++) {
		for (int j = 0; j < width; j++) {
			h_vec[i][j] = 1.f;
		}
	}

	// allocate 2D pitched device memory
	cudaMallocPitch(&dev_p, &pitch, width * sizeof(float), height);
	printf("pitch=%ld\n", pitch);

	cudaMemcpy2D(dev_p, pitch, h_vec, width *sizeof(float), width*sizeof(float), height, cudaMemcpyHostToDevice);

	dim3 grid_size(div_round_up(width, block_size_x), div_round_up(height, block_size_y));
	dim3 block_size(block_size_y, block_size_x);

	kernel_malloc_2d<<<grid_size, block_size>>>(dev_p, pitch, width, height);
	cudaDeviceSynchronize();

	cudaMemcpy2D(h_vec, width * sizeof(float), dev_p, pitch, width * sizeof(float), height, cudaMemcpyDeviceToHost);

	for (int i = 0; i < height; i++)  {
		for (int j = 0; j < width; j++)
			printf("row %i column %i value %f \n", i, j, h_vec[i][j]);
	}

	return 0;    
}
